// OpenCL Data Parallel Model (SIMD)
// Multiply Add: MADD d = (a x b) + c, equivalent to 2 flops, optional FMA.
// 64 Bit Floating Point Code.

#ifdef cl_amd_fp64
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
#endif
#ifdef cl_khr_fp64
#pragma OPENCL EXTENSION cl_khr_fp64: enable
#endif

#define madd(a, b, c) (a * b + c)

#define madd128(a, b) \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \

// 2 flops * 128 * 16 * 4 = 16384 flops
#define BENCHMARK \
    for (size_t i=0; i<4; i++) { \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
    } \

#if defined(cl_khr_fp64) || defined(cl_amd_fp64)

__kernel void fp64(__global double * const param0, __global double *param1) {
    const size_t gid = get_global_id(0);
    double a = (double)(param0[gid]);
    double b = (double)(1.0e-6);
    BENCHMARK
    param1[gid] = a + b;
}

__kernel void fp64_2(__global double * const param0, __global double *param1) {
    const size_t gid = get_global_id(0);
    double2 a = (double2)(param0[gid], param0[gid]);
    double2 b = (double2)(1.0e-6, 2.0e-6);
    BENCHMARK
    param1[gid] = (a.x + b.x) + (a.y + b.y);
}

__kernel void fp64_4(__global double * const param0, __global double *param1) {
    const size_t gid = get_global_id(0);
    double4 a = (double4)(param0[gid], param0[gid], param0[gid], param0[gid]);
    double4 b = (double4)(1.0e-6, 2.0e-6, 3.0e-6, 4.0e-6);
    BENCHMARK
    param1[gid] = (a.x + b.x) + (a.y + b.y) + (a.z + b.z) + (a.w + b.w);
}

__kernel void fp64_8(__global double * const param0, __global double *param1) {
    const size_t gid = get_global_id(0);
    double8 a = (double8)(param0[gid], param0[gid], param0[gid], param0[gid],
                          param0[gid], param0[gid], param0[gid], param0[gid]);
    double8 b = (double8)(1.0e-6, 2.0e-6, 3.0e-6, 4.0e-6,
                          5.0e-6, 6.0e-6, 7.0e-6, 8.0e-6);
    BENCHMARK
    param1[gid] = (a.s0 + b.s0) + (a.s1 + b.s1) + (a.s2 + b.s2) + (a.s3 + b.s3) +
                  (a.s4 + b.s4) + (a.s5 + b.s5) + (a.s6 + b.s6) + (a.s7 + b.s7);
}

/*__kernel void fp64_16(__global double * const param0, __global double *param1) {
    const size_t gid = get_global_id(0);
    double16 a = (double16)(param0[gid], param0[gid], param0[gid], param0[gid],
                            param0[gid], param0[gid], param0[gid], param0[gid],
                            param0[gid], param0[gid], param0[gid], param0[gid],
                            param0[gid], param0[gid], param0[gid], param0[gid]);
    double16 b = (double16)( 1.0e-6,  2.0e-6,  3.0e-6,  4.0e-6,
                             5.0e-6,  6.0e-6,  7.0e-6,  8.0e-6,
                             9.0e-6, 10.0e-6, 11.0e-6, 12.0e-6,
                            13.0e-6, 14.0e-6, 15.0e-6, 16.0e-6);
    BENCHMARK
    // NVidia GeForce 600M Graphics bug (v296).
    //param1[gid] = (a.s0 + b.s0) + (a.s1 + b.s1) + (a.s2 + b.s2) + (a.s3 + b.s3) +
    //              (a.s4 + b.s4) + (a.s5 + b.s5) + (a.s6 + b.s6) + (a.s7 + b.s7) +
    //              (a.s8 + b.s8) + (a.s9 + b.s9) + (a.sa + b.sa) + (a.sb + b.sb) +
    //              (a.sc + b.sc) + (a.sd + b.sd) + (a.se + b.se) + (a.sf + b.sf);
}*/

#endif
